#pragma once

#include <stdexcept>
#include <string>
#include <stdio.h>
#include <math.h>
#include "kernel_defs.hpp"
#include "maths_functions.hpp"


//template<gpu_size_t BLOCK_SIZE, typename Reall>
__global__
void pfb_device_kernel(const float* d_in,
                        float*      d_out,
                        float*         d_coeff,
                        gpu_size_t    count,
                        size_t        fft_size,
                        size_t        taps) 
{
  gpu_size_t thread_idx = threadIdx.x + blockIdx.x*blockDim.x;
  gpu_size_t thread_count = blockDim.x * gridDim.x;
   gpu_size_t add_idx = count + ((taps -1) * 2 * fft_size);
   int iIDx=0;
   int i;
   float c1,c2;
   gpu_size_t chan=0;    
// Iterate over values in the input
  for( gpu_size_t idx=thread_idx; idx<count; idx+=thread_count ) {
     chan = idx%(fft_size*2);
    for(i = 0; i < taps; ++i)
     {
          iIDx = (i * fft_size * 2);
      c1 +=  d_in[idx + iIDx] * d_coeff[iIDx + chan];  
      c2 +=  d_in[idx + add_idx + iIDx] * d_coeff[iIDx + chan]; 
      }
    d_out[idx] = c1;
    d_out[idx + count] = c2;  
   }
}




//template<typename Real>
void pfb_device(const float* d_in,
                 float*      d_out,
                 float*      d_coeff,
                 gpu_size_t    count,
                 size_t        fft_size,
                 size_t        taps,
                 cudaStream_t   stream = 0)
{
  enum { BLOCK_SIZE = 128  };
  dim3 block(BLOCK_SIZE);
  dim3 grid(math::min(math::div_round_up(count,
                                         (gpu_size_t)BLOCK_SIZE),
                      (gpu_size_t)CUDA_MAX_GRID_SIZE));
 // cout << "inside PFB kernel count= " << count << " fft_size= " << fft_size << " taps= " << taps  << " blockdim.x= " << block.x << " griddim.x= " << grid.x << endl; 
  pfb_device_kernel<<<grid, block, 0, stream>>>(d_in,
                                                d_out,
                                                d_coeff,
                                                count,
                                                fft_size,
                                                taps);
#ifdef _DEBUG
  cudaStreamSynchronize(stream);
  cudaError_t error = cudaGetLastError();
  if( error != cudaSuccess ) {
    throw std::runtime_error(std::string("In pfb_device_kernel: ") +
                             cudaGetErrorString(error));
  }
#endif
}

